/**
This file is shared by cuda-rawsha224 and cuda-rawsha256 formats, 
SHA256 definition is used to distinguish between them. 
*/
#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
#include <string.h>
#include "../cuda_rawsha256.h"

static void cuda_rawsha256(sha256_password *,void *);

#ifdef SHA256
  #define SHA_HASH sha256_hash
  __constant__ const uint32_t H[]={
   0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a, 0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19};
  extern "C" void gpu_rawsha256(sha256_password *i,SHA_HASH*o){cuda_rawsha256(i,o);}
#endif
#ifdef SHA224
  #define SHA_HASH sha224_hash
   __constant__ const uint32_t H[]={
   0xc1059ed8, 0x367cd507, 0x3070dd17, 0xf70e5939, 0xffc00b31,  0x68581511, 0x64f98fa7, 0xbefa4fa4};
   extern "C" void gpu_rawsha224(sha256_password *i,SHA_HASH*o){cuda_rawsha256(i,o);}
#endif

const uint32_t DATA_IN_SIZE=KEYS_PER_CRYPT*sizeof(sha256_password);
const uint32_t DATA_OUT_SIZE=KEYS_PER_CRYPT*sizeof(SHA_HASH);

static sha256_password *cuda_data=NULL;		///candidates
static SHA_HASH *cuda_data_out=NULL;		///sha256(candidate) or sha224(candidate)

__global__ void kernel_sha256(sha256_password *data,SHA_HASH *data_out);
static void cuda_rawsha256(sha256_password *host_in,void *out)
{
  SHA_HASH* host_out = (SHA_HASH *)out;
 ///Aloc memory and copy data to gpu
  cudaMalloc(&cuda_data,DATA_IN_SIZE);
  cudaMalloc(&cuda_data_out,DATA_OUT_SIZE);
  cudaMemcpy(cuda_data,host_in,DATA_IN_SIZE,cudaMemcpyHostToDevice);
 ///Run kernel and wait for execution end
  kernel_sha256<<<BLOCKS,THREADS>>>(cuda_data,cuda_data_out);
  cudaThreadSynchronize(); 
 ///Free memory and copy results back
  cudaMemcpy(host_out,cuda_data_out,DATA_OUT_SIZE,cudaMemcpyDeviceToHost);
  cudaFree(cuda_data);
  cudaFree(cuda_data_out);  
}

   
__global__ void kernel_sha256(sha256_password *data,SHA_HASH* data_out){ /// todo - use shared memory
  uint32_t idx = blockIdx.x*blockDim.x + threadIdx.x;
 const uint32_t k[]={
   0x428a2f98, 0x71374491, 0xb5c0fbcf, 0xe9b5dba5, 0x3956c25b, 0x59f111f1, 0x923f82a4, 0xab1c5ed5,
   0xd807aa98, 0x12835b01, 0x243185be, 0x550c7dc3, 0x72be5d74, 0x80deb1fe, 0x9bdc06a7, 0xc19bf174,
   0xe49b69c1, 0xefbe4786, 0x0fc19dc6, 0x240ca1cc, 0x2de92c6f, 0x4a7484aa, 0x5cb0a9dc, 0x76f988da,
   0x983e5152, 0xa831c66d, 0xb00327c8, 0xbf597fc7, 0xc6e00bf3, 0xd5a79147, 0x06ca6351, 0x14292967,
   0x27b70a85, 0x2e1b2138, 0x4d2c6dfc, 0x53380d13, 0x650a7354, 0x766a0abb, 0x81c2c92e, 0x92722c85,
   0xa2bfe8a1, 0xa81a664b, 0xc24b8b70, 0xc76c51a3, 0xd192e819, 0xd6990624, 0xf40e3585, 0x106aa070,
   0x19a4c116, 0x1e376c08, 0x2748774c, 0x34b0bcb5, 0x391c0cb3, 0x4ed8aa4a, 0x5b9cca4f, 0x682e6ff3,
   0x748f82ee, 0x78a5636f, 0x84c87814, 0x8cc70208, 0x90befffa, 0xa4506ceb, 0xbef9a3f7, 0xc67178f2};
  uint32_t w[64]={0};
  SHA_HASH* out=&data_out[idx];

  #pragma unroll 64
  for(uint32_t j=0;j<64;j++){
    if(j<16) w[j]=data[idx].v[j];
  else w[j]=sigma1(w[j-2])+w[j-7]+sigma0(w[j-15])+w[j-16];
  }
  
  uint32_t a=H[0];uint32_t b=H[1];uint32_t c=H[2];uint32_t d=H[3];
  uint32_t e=H[4];uint32_t f=H[5];uint32_t g=H[6];uint32_t h=H[7];
  #pragma unroll 64
  for(uint32_t j=0;j<64;j++){
   uint32_t t1=h+Sigma1(e)+Ch(e,f,g)+k[j]+w[j];
   uint32_t t2=Sigma0(a)+Maj(a,b,c);
   h=g;g=f;f=e;e=d+t1;d=c;c=b;b=a;a=t1+t2;
  }
  out->v[0]=a+H[0];out->v[1]=b+H[1];
  out->v[2]=c+H[2];out->v[3]=d+H[3];
  out->v[4]=e+H[4];out->v[5]=f+H[5];
  out->v[6]=g+H[6];
  #ifdef SHA256
  out->v[7]=h+H[7];
  #endif
}